(library (grid-gen OpenCL)
  (export OpenCL:make-code OpenCL:CPU)
  (import (guile)
          (rnrs)
          (grid-gen main)
          (grid-gen vars)
          (grid-gen code)
          (grid-gen symath)
          (grid-gen GPU)
          (except (srfi srfi-1) map)
  )

(define cur-indent (list "    "))
(define indent (lambda ()
  (apply string-append cur-indent)
))

(define OpenCL:CPU #f)

(define-syntax with-indent
  (syntax-rules ()
    ((_ v) (let ((old-ind cur-indent))
             (set! cur-indent (cons "  " cur-indent))
             (let ((r v))
               (set! cur-indent old-ind)
               r
             )
           )
    )
  )
)

(define float-type (lambda ()
  (cond ((= CODE:precision 4)
         "float"
        )
        ((= CODE:precision 8)
         "double"
        )
        (#t int-err "Invalid precision")
  )
))

(define var-type-s (lambda (vr)
  (if (is-int? vr)
      "int"
      (if (is-float? vr)
          (float-type)
          (int-err "Invalid var type")
      )
  )
))

(define cptr (list))

(define exp-ptrs (lambda (p)
  (map-ex (lambda (e) (if (is-ptr? e) (list-ref cptr (cadr e)) e))
          p
  )
))

(define edp (lambda (sp dp)
  (if (= CODE:precision 4) sp dp)
))

(define mk-num (lambda (n)
  (let* ((s (number->string n))
         (ss s)
        )
;    (if (and (not (string-has? (edp #\e #\d) ss)) (string-has? #\. ss))
;        (string-append ss "e0")
        ss
;    )
  )
))

;(define mk-num (lambda (n)
;  (let* ((s (number->string n))
;         (ss (edp s (string-map (lambda (x) (if (char-ci=? x #\e) #\d x)) s)))
;        )
;    (if (and (not (string-has? (edp #\e #\d) ss)) (string-has? #\. ss))
;        (string-append ss (edp "e0" "d0"))
;        ss
;    )
;  )
;))

(define mk-op (lambda (op args)
  (string-append "(" (string-join args op) ")")
))

(define mk-fn (lambda (fn args)
  (string-append fn "(" (string-join args ",") ")")
))

(define mk-fn-op (lambda (op args)
  (cond ((eq? op '+) (mk-op "+" args))
        ((eq? op '-) (mk-op "-" args))
        ((eq? op '*) (mk-op "*" args))
        ((eq? op '/) (mk-op "/" args))
        ((eq? op '**) (mk-fn "pow" args))
        ((eq? op 'sin) (mk-fn "sin" args))
        ((eq? op 'cos) (mk-fn "cos" args))
        ((eq? op 'tan) (mk-fn "tan" args))
        ((eq? op 'asin) (mk-fn "asin" args))
        ((eq? op 'acos) (mk-fn "acos" args))
        ((eq? op 'atan) (mk-fn "atan" args))
        ((eq? op 'log) (mk-fn "log" args))
        ((eq? op 'exp) (mk-fn "exp" args))
        ((eq? op 'sqrt) (mk-fn "sqrt" args))
        ((eq? op 'sign) (mk-fn "sign" args))
        ((eq? op 'abs) (mk-fn "fabs" args))
        ((eq? op 'min) (mk-fn "fmin" args))
        ((eq? op 'max) (mk-fn "fmax" args))
        ((eq? op 'pown) (mk-fn "pown" args))
        (#t (int-err "Unsupported op"))
  )
))

(define mk-ocl-src (lambda (lp cns-h red-h redf-h loc-h wsiz nam-h isiz osiz
                            in-pre in-loop out-pre out-loop
                           )
  (letrec*
    ((rnam-h (make-hash-table))
    
     (mk-vr-hs (lambda (l)
       (let ((h (make-hash-table)))
         (for-each
           (lambda (iv)
             (hash-set! h (simp-ex (cadr iv)) (* (car iv) wsiz))
           )
           (zip (iota1 (length l)) l)
         )
         h
       )
     ))
     
     (in-pre-h (mk-vr-hs in-pre))
     (in-loop-h (mk-vr-hs in-loop))
     (out-pre-h (mk-vr-hs out-pre))
     (out-loop-h (mk-vr-hs out-loop))
     
     (cur-vr-h-in #f)
     (cur-vr-h-out #f)
     
     (g-f-name (lambda (n i)
       (let ((n1 (string-append n (if (> i 0) (number->string i) ""))))
         (if (hash-ref rnam-h n1 #f)
             (g-f-name n (+ i 1))
             n1
         )
       )
     ))
     
     (var-full-name (lambda (v)
       (let ((nm (hash-ref nam-h (var-id v) #f)))
         (if nm
             nm
             (let ((nm1 (g-f-name (var-name v) 0)))
               (hash-set! nam-h (var-id v) nm1)
               (hash-set! rnam-h nm1 #t)
               nm1
             )
         )
       )
     ))
     
     (rp-mm (lambda (f)
       (if (= (length f) 3)
           (if (equal? (cadr f) (caddr f))
               (cadr f)
               f
           )
           (if (= (length f) 2)
               (cadr f)
               (if (> (length f) 3)
                   (list (car f)
                         (cadr f)
                         (rp-mm (cons (car f) (cddr f)))
                   )
                   f
               )
           )
           
       )
     ))
     
     (mk-ex (lambda (ex)
       (cond ((number? ex)
              (if (< ex 0)
                  (string-append "(" (mk-num ex) ")")
                  (mk-num ex)
              )
             )
             ((is-var? ex)
              (if (hash-ref loc-h (var-id ex) #f)
                  (var-full-name ex)
                  (let ((cn (hash-ref cns-h ex #f)))
                    (if cn
                        (format #f "cnst~a" cn)
                        (let ((rn (hash-ref red-h (var-id ex) #f)))
                          (if rn
                              (format #f "(l_trd~a)" rn)
                              (let ((res (format #f "ibuf[ipt+~a]"
                                            (hash-ref cur-vr-h-in ex #f)
                                         )
                                    )
                                   )
                                res
                              )
                          )
                        )
                    )
                  )
              )
             )
             ((is-ptr? ex)
              (list-ref cptr (cadr ex))
             )
             ((or (is-op? ex) (is-fn? ex))
              (let ((ex2 (cond ((and (eq? (car ex) '**) (number? (caddr ex)))
                                (cond ((= (caddr ex) 0.5)
                                      `(sqrt ,(cadr ex))
                                      )
                                      ((= (caddr ex) 2)
                                      `(* ,(cadr ex) ,(cadr ex))
                                      )
                                      ((= (floor (caddr ex)) (caddr ex))
                                       `(pown ,(cadr ex) 
                                              ,(inexact->exact (caddr ex))
                                        )
                                      )
                                      (#t ex)
                                )
                               )
                               ((or (eq? (car ex) 'min)
                                    (eq? (car ex) 'max)
                                )
                                (let ((rt (rp-mm ex)))
                                  rt
                                )
                               )
                               (#t ex)
                         )
                    )
                   )
                (if (or (is-op? ex2) (is-fn? ex2))
                    (mk-fn-op (car ex2) (map mk-ex (cdr ex2)))
                    (mk-ex ex2)
                )
              )
             )
             (#t  (format #t "EXPR: ~a\n" ex)
                  (int-err "Unsupported expr")
             )
       )
     ))
     (mk-for (lambda (dms seq cod pre)
       (set! cptr
             (map (lambda (i)
                    (let ((nm (format #f "i~a" i)))
                      (hash-set! rnam-h nm #t)
                      nm
                    )
                  )
                  (iota1 (length dms))
             )
       )
       (string-append
          (if (> isiz 0)
              (string-append (indent) "int ipt=myid;\n")
              ""
          )
          (if (> osiz 0)
              (string-append (indent) "int opt=myid;\n")
              ""
          )
          (apply string-append
            (map (lambda (i)
                   (format #f "~a~a i~a~a;\n"
                     (indent)
                     (float-type)
                     i
                     (if (> i 0)
                         (format #f "=ibuf[ipt+~a]" (* wsiz (- i 1)))
                         ""
                     )
                   )
                 )
                 (iota1 (length seq))
            )
          )
          (indent) (format #f "ipt+=~a;\n" (* wsiz (- (length seq) 1)))
          (if pre
              (let ((optr cptr))
                 (set! cptr
                       (cons (mk-ex (simp-ex (car (list-ref dms (car seq)))))
                             (cdr cptr)
                       )
                 )
                 (set! cur-vr-h-in in-pre-h)
                 (set! cur-vr-h-out out-pre-h)
                 (let* ((res (mk-cd pre)))
                   (set! cptr optr)
                   (string-append
                     res
                     (if (not (null? in-pre))
                         (format #f "~aipt+=~a;\n" 
                           (indent) (* (length in-pre) wsiz)
                         )
                         ""
                     )
                     (if (not (null? out-pre))
                         (format #f "~aopt+=~a;\n"
                           (indent) (* (length out-pre) wsiz)
                         )
                         ""
                     )
                   )
                 )
              )
              ""
          )
          (indent) (format #f "for(i0=~a;i0<~a;i0+=1.0) {\n"
                     (mk-ex (simp-ex (car  (list-ref dms (car seq)))))
                     (mk-ex (simp-ex `(+ ,(cadr (list-ref dms (car seq))) 1)))
                   )
          (begin
            (set! cur-vr-h-in in-loop-h)
            (set! cur-vr-h-out out-loop-h)
            (with-indent
              (string-append
                (mk-cd cod)
                (if (not (null? in-loop))
                    (format #f "~aipt+=~a;\n"
                      (indent) (* (length in-loop) wsiz)
                    )
                    ""
                )
                (if (not (null? out-loop))
                    (format #f "~aopt+=~a;\n"
                      (indent) (* (length out-loop) wsiz)
                    )
                    ""
                )
              )
            )
          )
          (indent) "}\n"
       )
     ))
     (mk-sv (lambda (v0 e0)
       (let ((vr (if (hash-ref loc-h (var-id v0) #f)
                     (var-full-name v0)
                     (let ((rv (hash-ref red-h (var-id v0) #f)))
                       (if rv
                           (format #f "l_trd~a" rv)
                           (let ((res (format #f "obuf[opt+~a]"
                                         (hash-ref cur-vr-h-out (simp-ex v0) #f)
                                      )
                                 )
                                )
                             res
                           )
                       )
                     )
                 )
             )
             (ex (mk-ex (simp-ex e0)))
            )
         (string-append (indent) vr "=" ex ";\n")
       )
     ))
     (mk-cd (lambda (cd)
       (let ((c (car cd)))
         (cond ((eq? c 'CODE_BEGIN)
                (apply string-append (map mk-cd (cdr cd)))
               )
               ((eq? c 'OMP_DO) "")
               ((eq? c 'OMP_REDUCE) "")
               ((eq? c 'OMP_LOCALS) "")
               ((eq? c 'CODE_NOP) "")
               ((eq? c 'RAW) (format #f "//~a\n" (cadr cd)))
               ((eq? c 'SET_PTR) 
                (set! cptr (map mk-ex (map simp-ex (cadr cd))))
                ""
               )
               ((eq? c 'CODE_LOOP)
                (mk-for (cadr (obj-find-el cd 'LOOP_DIMS))
                        (cadr (obj-find-el cd 'LOOP_SEQ))
                        (cadr (obj-find-el cd 'LOOP_CODE))
                        (let ((pc (obj-find-el cd 'LOOP_PRE_CODE)))
                          (if pc (cadr pc) #f)
                        )
                )
               )
               ((eq? c 'SET_VAL)
                (mk-sv (cadr cd) (caddr cd))
               )
               (#t (begin
                     (format #t "CM: ~a\n" c)
                     (int-err "Unsupported code")
                   )
               )
         )
       )
     ))
    )
    (mk-cd lp)
  )
))

(define make-kernel (lambda (kid in-pre in-loop out-pre out-loop
                             cnstl loc-vars reds lp wsiz isiz osiz
                            )
  (string-append
    (format #f "__kernel void kernel_~a(~a) {\n"
       kid
       (string-join
         (append
           (if (> isiz 0)
               (list (format #f "__global ~a *ibuf" (float-type)))
              '()
           )
           (if (> osiz 0)
               (list (format #f "__global ~a *obuf" (float-type)))
              '()
           )
           (map (lambda (iv)
                  (format #f "const ~a cnst~a"
                    (var-type-s (cadr iv))
                    (car iv)
                  )
                )
                (zip (iota1 (length cnstl)) cnstl)
           )
           (map (lambda (iv)
                  (format #f "__global ~a *rdv~a,__local ~a *trd~a"
                    (var-type-s (cadr iv))
                    (car iv)
                    (var-type-s (cadr iv))
                    (car iv)
                  )
                )
                (zip (iota1 (length reds)) (map car (map reverse reds)))
           )
         )
         ","
       )
    )
    (let ((cns-h (make-hash-table))
          (red-h (make-hash-table))
          (redf-h (make-hash-table))
          (loc-h (make-hash-table))
          (nam-h (make-hash-table))
         )
      (for-each
        (lambda (iv) (hash-set! cns-h (cadr iv) (car iv)))
        (zip (iota1 (length cnstl)) cnstl)
      )
      (for-each
        (lambda (iv)
          (let ((v (car (reverse (cadr iv))))
                (fn (cadr (obj-find-el (cadr iv) 'OMP_FUNC)))
                (i (car iv))
               )
            (hash-set! redf-h i fn)
            (hash-set! red-h (var-id v) i)
          )
        )
        (zip (iota1 (length reds)) reds)
      )
      (for-each
        (lambda (v) (hash-set! loc-h (var-id v) v))
        loc-vars
      )
      (letrec* ((res (mk-ocl-src
                       lp cns-h red-h redf-h loc-h wsiz nam-h isiz osiz
                       in-pre in-loop out-pre out-loop
                     )
                )
                (lvrs (hash-map->list (lambda (k v) v) loc-h))
                (iloc (filter is-int? lvrs))
                (floc (filter is-float? lvrs))
                (chop-lst (lambda (l)
                  (if (< (length l) 10)
                      (list l)
                      (cons
                        (list-head l 10)
                        (chop-lst (list-tail l 10))
                      )
                  )
                ))
            )
        (string-append
          (if (null? iloc)
              ""
              (string-append
                (indent) "int "
                (string-join
                  (filter
                    (lambda (x) x)
                    (map (lambda (v)
                           (hash-ref nam-h (var-id v) #f)
                         )
                         iloc
                    )
                  )
                  ","
                )
                ";\n"
              )
          )
          (if (null? floc)
              ""
              (apply string-append
                (map
                  (lambda (l)
                    (string-append
                      (indent) (float-type) " "
                      (string-join l ",")
                      ";\n"
                    )
                  )
                  (chop-lst
                    (filter
                      (lambda (x) x) 
                      (map (lambda (v) (hash-ref nam-h (var-id v) #f)) floc)
                    )
                  )
                )
              )
          )
          (apply string-append
            (map
              (lambda (iv)
                (format #f "~a~a l_trd~a=rdv~a[0];\n"
                  (indent)
                  (var-type-s (cadr iv))
                  (car iv)
                  (car iv)
                )
              )
              (zip (iota1 (length reds)) (map car (map reverse reds)))
            )
          )
          (indent) "int myid=get_global_id(0);\n"
          res
          (apply string-append
            (map
              (lambda (ird)
                (letrec* 
                  ((i (car ird))
                   (rd (cadr ird))
                   (vr (car (reverse rd)))
                   (fn (cadr (obj-find-el rd 'OMP_FUNC)))
                   (st (inexact->exact (floor (/ (log wsiz) (log 2)))))
                   (b1 (expt 2 st))
                   (bnd (if (< b1 wsiz) b1 (/ b1 2)))
                   (mk-rd (lambda (b)
                     (if (= b 1)
                         (format #f "~aif(myid==0) rdv~a[get_group_id(0)]=~a;\n"
                           (indent)
                           i
                           (mk-fn-op
                             fn
                             (list (format #f "trd~a[0]" i)
                                   (format #f "trd~a[1]" i)
                             )
                           )
                         )
                         (string-append
                           (format #f 
                             "~aif(myid<~a) trd~a[myid]=~a;\n"
                             (indent)
                             b
                             i
                             (let ((op (mk-fn-op
                                         fn 
                                         (list (format #f "trd~a[myid]" i)
                                               (format #f "trd~a[myid+~a]" 
                                                 i
                                                 b
                                               )
                                         )
                                       )
                                   )
                                  )
                               (format #f "(myid+~a<loc_size)?~a:l_trd~a"
                                     b op i
                               )
                             )
                           )
                           (indent) "barrier(CLK_LOCAL_MEM_FENCE);\n"
                           (mk-rd (inexact->exact (floor (/ b 2))))
                         )
                     )
                   ))
                  )
                  (string-append
                    (indent) "int loc_size=get_local_size(0);\n"
                    (indent) "myid=get_local_id(0);\n"
                    (indent) (format #f "trd~a[myid]=l_trd~a;\n" i i)
                    (indent) "barrier(CLK_LOCAL_MEM_FENCE);\n"
                    (mk-rd bnd)
                  )
                )
              )
              (zip (iota1 (length reds)) reds)
            )
          )
        )
      )
    )
    "}\n\n"
  )
))

(define make-kernels (lambda (krns)
  (apply string-append
    (hash-map->list
      (lambda (kid krl)
        (let ((in-pre   (list-ref krl 0))
              (in-loop  (list-ref krl 1))
              (out-pre  (list-ref krl 2))
              (out-loop (list-ref krl 3))
              (cnstl    (list-ref krl 4))
              (loc-vars (list-ref krl 5))
              (reds     (list-ref krl 6))
              (lp       (list-ref krl 7))
              (wsiz     (list-ref krl 8))
              (isiz     (list-ref krl 9))
              (osiz     (list-ref krl 10))
             )
          (make-kernel kid in-pre in-loop out-pre out-loop
                       cnstl loc-vars reds lp wsiz isiz osiz
          )
        )
      )
      krns
    )
  )
))

(define make-quotes (lambda (str)
  (apply string-append
    (map
      (lambda (s)
        (string-append "\"" s "\\n\"\n")
      )
      (string-split str #\newline)
    )
  )
))

(define mk-code (lambda (krns)
  (letrec* ((nk (length (hash-map->list (lambda (k v) 1) krns)))
           )
    (string-append
      "#include <stdio.h>\n"
      "#include <OpenCL/opencl.h>\n"
      "#include <math.h>\n"
      "\n"
      "#define MAX_DEV 32\n"
      "#define ERR(er) {fprintf(stderr,\"Error: %s\\n\",er);abort();}\n"
      "\n"
      "cl_uint ndev;\n"
      "cl_device_id device_id[MAX_DEV];\n"
      "cl_context context;\n"
      "cl_command_queue queue;\n"
      "cl_program program;\n"
      "\n"
      "char *KernelSource = \n"
        (if (= CODE:precision 8)
            "\"#pragma OPENCL EXTENSION cl_khr_fp64 : enable\\n\"\n"
            ""
        )
        (make-quotes (make-kernels krns))
      ";\n\n"
      (format #f "cl_kernel kernel[~a];\n" nk)
      (format #f "cl_mem out_args[~a];\n" nk)
      (format #f "size_t kern_loc[~a];\n" nk)
      (format #f "int kern_grp[~a];\n" nk)
      (apply string-append
        (map (lambda (kern-id)
               (let* ((krl (hash-ref krns kern-id #f))
                      (in-pre   (list-ref krl 0))
                      (in-loop  (list-ref krl 1))
                      (out-pre  (list-ref krl 2))
                      (out-loop (list-ref krl 3))
                      (cnstl    (list-ref krl 4))
                      (loc-vars (list-ref krl 5))
                      (reds     (list-ref krl 6))
                      (lp       (list-ref krl 7))
                      (wsiz     (list-ref krl 8))
                      (isiz     (list-ref krl 9))
                      (osiz     (list-ref krl 10))
                     )
                 (if (not (null? reds))
                     (format #f "cl_mem kern~a_rdb[~a];\n"
                        kern-id (length reds)
                     )
                     ""
                 )
               )
             )
             (iota1 nk)
        )
      )
      "\n"
   ;   "char *rep_str(char *src, char *s1, char *s2) {\n"
   ;   "  char *res=malloc(1);\n"
   ;   "  char *t;\n"
   ;   "  int l1=strlen(s1);\n"
   ;   "  int l2=strlen(s2);\n"
   ;   "  int lt;\n"
   ;   "  int lr=0;\n"
   ;   "  res[0]='\\0';\n"
   ;   "  while((t=strstr(src,s1))!=NULL) {\n"
   ;   "    lt=strlen(t);\n"
   ;   "    res=realloc(lr+l3+l2);\n"
   ;   "    memcpy(res+lr,src,lt);\n"
   ;   "    memcpy(res+lr+lt,s2,l2);\n"
   ;   "    lr+=lt+l2;\n"
   ;   "    res[lr]='\\0';\n"
   ;   "    src+=lt;\n"
   ;   "  }\n"
   ;   "  lt=strlen(src);\n"
   ;   "  res=realloc(lr+lt);\n"
   ;   "  memcpy(res+lr,src,lt);\n"
   ;   "  res[lr+lt]='\\0';\n"
   ;   "  return(res);\n"
   ;   "}\n"
      "\n"
      "void gpu_init_(void) {\n"
      (if (= CODE:precision 8)
          "  cl_device_fp_config fpc;\n"
          ""
      )
      "  cl_uint max_consts;\n"
      "  char dev_name[1024];\n"
      "  int err;\n"
      "  char err_buf[10240];\n"
      "  size_t err_len;\n"
      "  size_t tlc;\n"
      "  int dv;\n"
      "  char tnum[128];\n"
      "\n"
      "  if(clGetDeviceIDs(NULL,CL_DEVICE_TYPE_"(if OpenCL:CPU "CPU" "GPU")
           ",MAX_DEV,device_id,&ndev)!=CL_SUCCESS)\n"
      "    ERR(\"clGetDeviceIDs error\");\n"
      "  if(ndev<=0) ERR(\"No OpenCL compatible devices found!\");\n"
      "  printf(\"Found %d OpenCL devices.\\n\", ndev);\n"
      "  if(clGetDeviceInfo(device_id[0],CL_DEVICE_MAX_CONSTANT_ARGS,sizeof(max_consts),&max_consts,NULL)!=CL_SUCCESS)\n"
      "    ERR(\"clGetDeviceInfo(CL_DEVICE_MAX_CONSTANT_ARGS) error\");\n"
      "//  if(max_consts<" (number->string GPU:max-consts) ") ERR(\"This code required too many GPU consts\");\n"
      "  if(clGetDeviceInfo(device_id[0],CL_DEVICE_NAME,1024,dev_name,NULL)!=CL_SUCCESS)\n"
      "    ERR(\"clGetDeviceInfo(CL_DEVICE_NAME) error\");\n"
      "  printf(\"Device name: %s\\n\",dev_name);\n"
      (if (= CODE:precision 8)
          (string-append
            "  if(clGetDeviceInfo(device_id[0],CL_DEVICE_DOUBLE_FP_CONFIG,sizeof(fpc),&fpc,NULL)!=CL_SUCCESS)\n"
            "    ERR(\"clGetDeviceInfo(CL_DEVICE_DOUBLE_FP_CONFIG) error\");\n"
            "  if(fpc==0) ERR(\"Double precision is not supported by GPU device\");\n"
          )
          ""
      )
      "  context=clCreateContext(0,1,&device_id[0],NULL,NULL,&err);\n"
      "  if(!context) ERR(\"Cannot create context.\");\n"
      "  queue=clCreateCommandQueue(context,device_id[0],0,&err);\n"
      "  if(!queue) ERR(\"Cannot create queue.\");\n"
      "  program=clCreateProgramWithSource(context,1,(const char **)&KernelSource,NULL,&err);\n"
      "  if(!program) ERR(\"Cannot create program.\");\n"
      "  int comp_err=clBuildProgram(program,1,device_id,\""
            (if (= CODE:precision 4)
                " -cl-single-precision-constant"
                ""
            )
          "\",NULL,NULL);\n"
      "  clGetProgramBuildInfo(program,device_id[0],CL_PROGRAM_BUILD_LOG,sizeof(err_buf),err_buf,&err_len);\n"
      "  fprintf(stderr,\"Kernel compile output:\\n%s\\n\", err_buf);\n"
      "  if(comp_err!=CL_SUCCESS) ERR(\"Kernel compile error.\");\n"
      (apply string-append
             (map (lambda (i)
                    (string-append
      (format #f "  kernel[~a]=clCreateKernel(program,\"kernel_~a\",&err);\n" i i)
      (format #f "  if(!kernel[~a]||err!=CL_SUCCESS) ERR(\"Cannot create kernel ~a.\")\n" i i)
      (format #f "  if(clGetKernelWorkGroupInfo(kernel[~a],device_id[0],CL_KERNEL_WORK_GROUP_SIZE,sizeof(tlc),&tlc,NULL)!=CL_SUCCESS)\n" i)
                 "    ERR(\"Cannot get local group size for kernel "(number->string i)"\");\n"
               (let* ((krl  (hash-ref krns i #f))
                      (wsiz (list-ref krl 8))
                     )
                 (string-append
                   "  for(dv=1;dv<="(number->string wsiz)";dv++) {\n"
                   "    div_t dt=div("(number->string wsiz)",dv);\n"
                   "    if(dt.rem==0&&dt.quot<=tlc) {\n"
                   "      kern_loc["(number->string i)"]=dt.quot;\n"
                   "      kern_grp["(number->string i)"]=dv;\n"
                   "      printf(\"Local size for kernel "(number->string i)": %d\\n\",dt.quot);\n"
                   "      snprintf(tnum,128,\"%d\",dt.quot);\n"
                   "      "
                   "      break;\n"
                   "    }\n"
                   "  }\n"
                 )
               )
                    )
                  )
                  (iota1 nk)
             )
      )
      "}\n"
      "\n"
      "void gpu_finish_(void) {\n"
      "  int i;\n"
      "  for(i=0;i<"(number->string nk)";i++) {\n"
      "    if(out_args[i]) {\n"
      "      clReleaseMemObject(out_args[i]);\n"
      "      out_args[i]=NULL;\n"
      "    }\n"
      "  }\n"
      (apply string-append
        (map
          (lambda (i)
            (format #f "  clReleaseKernel(kernel[~a]);\n" i)
          )
          (iota1 nk)
        )
      )
      "  clReleaseProgram(program);\n"
      "  clReleaseCommandQueue(queue);\n"
      "  clReleaseContext(context);\n"
      "}\n"
      "\n"
      (apply
        string-append
        (map 
          (lambda (kern-id)
            (let* ((krl (hash-ref krns kern-id #f))
                   (in-pre   (list-ref krl 0))
                   (in-loop  (list-ref krl 1))
                   (out-pre  (list-ref krl 2))
                   (out-loop (list-ref krl 3))
                   (cnstl    (list-ref krl 4))
                   (loc-vars (list-ref krl 5))
                   (reds     (list-ref krl 6))
                   (lp       (list-ref krl 7))
                   (wsiz     (list-ref krl 8))
                   (isiz     (list-ref krl 9))
                   (osiz     (list-ref krl 10))
                   
                   (reds-v   (map car (map reverse reds)))
                   (inb-cnt  (if (> isiz 0) 1 0))
                   (outb-cnt  (if (> osiz 0) 1 0))
                  )
              (string-append
                (format #f "void gpu_run_kernel~a_(~a) {\n" 
                  kern-id
                  (string-join
                    (append
                      (if (> isiz 0)
                          (list
                            (string-append (float-type) " *ibuf")
                          )
                         '()
                      )
                      (map
                        (lambda (iv)
                          (format #f "~a *cnst~a"
                            (var-type-s (cadr iv))
                            (car iv)
                          )
                        )
                        (zip (iota1 (length cnstl)) cnstl)
                      )
                      (map (lambda (iv)
                             (format #f "~a *rdv~a" 
                                 (var-type-s (cadr iv)) (car iv)
                             )
                           )
                           (zip (iota1 (length reds-v)) reds-v)
                      )
                    )
                    ","
                  )
                )
                
                (if (> isiz 0)
                    (string-append
                      (format #f "  cl_mem in_bf=clCreateBuffer(context,CL_MEM_READ_ONLY,sizeof(~a)*~a,NULL,NULL);\n" 
                         (float-type) isiz)
                      (format #f "  if(!in_bf) ERR(\"Cannot allocate input buffer (kernel ~a).\");\n" kern-id)
                      (format #f "  if(clEnqueueWriteBuffer(queue,in_bf,CL_TRUE,0,sizeof(~a)*~a,ibuf,0,NULL,NULL)!=CL_SUCCESS)\n" (float-type) isiz)
                      "    ERR(\"Cannot copy input data to device\");\n"
                      "  if(clSetKernelArg(kernel["(number->string kern-id)"],0,sizeof(cl_mem),&in_bf)!=CL_SUCCESS)\n"
                      (format #f "    ERR(\"Cannot set kernel arg (0, kernel ~a)\");\n" kern-id)
                    )
                    ""
                )

                (if (> osiz 0)
                    (string-append
                      (format #f "  out_args[~a]=clCreateBuffer(context,CL_MEM_WRITE_ONLY,sizeof(~a)*~a,NULL,NULL);\n" kern-id (float-type) osiz)
                      (format #f "  if(!out_args[~a]) ERR(\"Cannot allocate output buffer (kernel ~a).\");\n" kern-id kern-id)
                      (format #f "  if(clSetKernelArg(kernel[~a],~a,sizeof(cl_mem),&out_args[~a])!=CL_SUCCESS)\n" kern-id inb-cnt kern-id)
                      (format #f "    ERR(\"Cannot set kernel arg (~a, kernel ~a)\");\n" inb-cnt kern-id)
                    )
                    ""
                )
                                
                (apply string-append
                  (map (lambda (iv)
                          (string-append
                             "  if(clSetKernelArg(kernel["(number->string kern-id)"],"
                                     (number->string (+ (car iv) inb-cnt outb-cnt))",sizeof("
                                     (var-type-s (cadr iv))
                                     "),cnst"(number->string (car iv))")!=CL_SUCCESS)\n"
                             (format #f "    ERR(\"Cannot set kernel const arg (~a, kernel ~a)\");\n" (car iv) kern-id)
                          )
                       )
                       (zip (iota1 (length cnstl)) cnstl)
                  )
                )
                (apply string-append
                  (map (lambda (iv)
                          (string-append
                            (format #f "  kern~a_rdb[~a]=clCreateBuffer(context,CL_MEM_READ_WRITE,sizeof(~a)*kern_grp[~a],NULL,NULL);\n"
                              kern-id (car iv)
                              (var-type-s (cadr iv))
                              kern-id
                            )
                            (format #f "  if(!kern~a_rdb[~a]) ERR(\"Cannot allocate buffer for reduce var ~a for kernel ~a\");\n"
                              (car iv) kern-id (car iv) kern-id
                            )
                            (format #f "  if(clEnqueueWriteBuffer(queue,kern~a_rdb[~a],CL_TRUE,0,sizeof(~a),rdv~a,0,NULL,NULL)!=CL_SUCCESS)\n"
                               kern-id (car iv) (var-type-s (cadr iv)) (car iv)
                            )
                            (format #f "    ERR(\"Cannot copy input rdv ~a to device\");\n" (car iv))
                            (format #f "  if(clSetKernelArg(kernel[~a],~a,sizeof(cl_mem),&kern~a_rdb[~a])!=CL_SUCCESS)\n"
                                     kern-id
                                     (+ (* 2 (car iv)) (length cnstl) inb-cnt outb-cnt)
                                     kern-id
                                    (car iv)
                            )
                            (format #f "    ERR(\"Cannot set kernel ~a rdb arg ~a\");\n" kern-id (car iv))
                            (format #f "  if(clSetKernelArg(kernel[~a],~a,sizeof(~a)*kern_loc[~a],NULL)!=CL_SUCCESS)\n"
                                     kern-id
                                     (+ (* 2 (car iv)) 1 (length cnstl) inb-cnt outb-cnt)
                                     (var-type-s (cadr iv))
                                     kern-id
                            )
                            (format #f "    ERR(\"Cannot set kernel ~a trd arg ~a\");\n" kern-id (car iv))
                          )
                       )
                       (zip (iota1 (length reds-v)) reds-v)
                  )
                )
                "  size_t glb[1]={" (number->string wsiz) "};\n"
                "  size_t loc[1]={kern_loc["(number->string kern-id)"]};\n"
                (format #f "  if(clEnqueueNDRangeKernel(queue,kernel[~a],1,NULL,glb,loc,0,NULL,NULL)!=CL_SUCCESS)\n" kern-id)
                (format #f "    ERR(\"Cannot start kernel ~a.\");\n" kern-id)
                
                "  clFlush(queue);\n"
                "  clFinish(queue);\n"

                (format #f "  clReleaseMemObject(in_bf);\n")
                "}\n"
                "\n"
                (if (or (not (null? reds)) (> osiz 0))
                    (string-append
                      (format #f "void gpu_get_result~a_(~a) {\n"
                          kern-id
                          (string-join
                            (append
                              (if (> osiz 0)
                                  (list
                                    (format #f "~a *obuf" (float-type))
                                  )
                                 '()
                              )
                              (map
                                (lambda (rd)
                                  (let ((vr (car (reverse (cadr rd)))))
                                    (format #f "~a *r~a"
                                      (var-type-s vr)
                                      (car rd)
                                    )
                                  )
                                )
                                (zip (iota1 (length reds)) reds)
                              )
                            )
                            ","
                          )
                      )
                      
                      (if (> osiz 0)
                          (string-append
                            (format #f "  if(clEnqueueReadBuffer(queue,out_args[~a],CL_TRUE,0,sizeof(~a)*~a,obuf,0,NULL,NULL)!=CL_SUCCESS)\n"
                              kern-id (float-type) osiz
                            )
                            "    ERR(\"Cannot copy output data from device\");\n"
                            (format #f "  clReleaseMemObject(out_args[~a]);\n" kern-id)
                            (format #f "  out_args[~a]=NULL;\n" kern-id)
                          )
                          ""
                      )
                      (apply string-append
                        (map (lambda (iv)
                               (let ((i (car iv))
                                     (v (car (reverse (cadr iv))))
                                     (f (cadr (obj-find-el (cadr iv) 'OMP_FUNC)))
                                    )
                                 (string-append
                                    (format #f "  ~a *trdb~a;\n"
                                      (var-type-s v)
                                      i
                                    )
                                    (format #f "  trdb~a=malloc(kern_grp[~a]*sizeof(~a));\n"
                                      i
                                      kern-id
                                      (var-type-s v)
                                    )
                                    (format #f "  if(clEnqueueReadBuffer(queue,kern~a_rdb[~a],CL_TRUE,0,sizeof(~a)*kern_grp[~a],trdb~a,0,NULL,NULL)!=CL_SUCCESS)\n"
                                            kern-id
                                            i
                                            (var-type-s v)
                                            kern-id
                                            i
                                    )
                                    (format #f "    ERR(\"Cannot get kernel ~a rdb ~a\");\n" kern-id i)
                                    (format #f "  clReleaseMemObject(kern~a_rdb[~a]);\n" kern-id i)
                                    (format #f "  kern~a_rdb[~a]=NULL;\n" kern-id i)
                                    (format #f "  *r~a=trdb~a[0];\n" i i)
                                    

                                    "  int i;\n"
                                    
                                    (format #f "  for(i=1;i<kern_grp[~a];i++)\n" kern-id)
                                    (format #f "    *r~a=~a;\n"
                                      i
                                      (mk-fn-op
                                        f
                                        (list 
                                          (format #f "*r~a" i)
                                          (format #f "trdb~a[i]" i)
                                        )
                                      )
                                    )
                                   
                                    (format #f "  free(trdb~a);\n" i)
                                 )
                               )
                             )
                             (zip (iota1 (length reds)) reds)
                        )
                      )
                      "}\n"
                      "\n"
                    )
                    ""
                )
                "\n"
              )
            )
          )
          (iota1 nk)
        )
      )
    )
  )
))

(define OpenCL:make-code (lambda (krns)
  (letrec* ((f (transcoded-port
                 (open-file-output-port
                   "opencl_code.c"
                   (file-options no-fail)
                 )
                 (native-transcoder)
               )
            )
            (cod (mk-code krns))
           )
    (format f "~a" cod)
    (close-port f)
  )

))  

)

(set! GPU:make-code OpenCL:make-code)

